#include "CL/cl.h"
#include <sys/types.h>
#include <stdio.h>
#include <string.h>
#include <sys/time.h>

#include <clAmdBlas.h>

#define M (2048)
#define N (2048)
#define K (2048)
#define concur (16)
#define TILE (16)

static const clAmdBlasOrder order = clAmdBlasRowMajor;
static const clAmdBlasTranspose transA = clAmdBlasNoTrans;
static const clAmdBlasTranspose transB = clAmdBlasNoTrans;

# define timersub(a, b, result)                                         \
        do {                                                            \
                (result)->tv_sec = (a)->tv_sec - (b)->tv_sec;           \
                (result)->tv_usec = (a)->tv_usec - (b)->tv_usec;        \
                if ((result)->tv_usec < 0) {                            \
                        --(result)->tv_sec;                             \
                        (result)->tv_usec += 1000000;                   \
                }                                                       \
        } while (0)

const char *KernelSource3 = "\n"                                                         \
"__kernel void kernel3(__global float *A, __global float *B, __global float *C          \n"\
"					   int m, int n, int k)												\n"\
"{                                                                                      \n"\
"       const int TILE =16;                                                           \n"\
"       int idx = get_local_id(0);                                                      \n"\
"       int idy = get_local_id(1);                                                      \n"\
"       int row = (get_group_id(1)*TILE)+idy;                                           \n"\
"       int col = (get_group_id(0)*TILE)+idx;                                           \n"\
"       int i;                                                                          \n"\
"       int x;                                                                          \n"\
"       __local float Ab[TILE][TILE];                                                       \n"\
"       __local float Bb[TILE][TILE];                                                       \n"\
"       float sum = 0;                                                                  \n"\
"       for(i=0 ; i<k ; i+=TILE){                                                       \n"\
"               Ab[idy][idx]=A[row*k+(idx+i)];                                          \n"\
"               Bb[idy][idx]=B[(idy+i)*n+col];                                          \n"\
"               barrier(CLK_LOCAL_MEM_FENCE);                                           \n"\
"               for(x = 0 ; x < TILE ; ++x){                                            \n"\
"                       sum = sum + (Ab[idy][x] * Bb[x][idx]);                          \n"\
"               }                                                                       \n"\
"				barrier(CLK_LOCAL_MEM_FENCE);											\n"\
"       }                                                                               \n"\
"       C[row * n + col] = sum;                                                         \n"\
"                                                                                       \n"\
"                                                                                       \n"\
"}                                                                                      \n"\
"\n";

size_t shrRoundUp(int group_size, int global_size)
{
	int r = global_size % group_size;
	if(r == 0)
	{
		return global_size;
	}
	 else
	{
		return global_size + group_size - r;
	}
}

int main()
{
	struct timeval time_start, time_end, time_res;

	int count = 0;
	float alpha = 1.0;
	float beta = 0.0;

	float *m1_h, *m2_h, *o_h;
	cl_mem m1_d,m2_d,o_d;
	float *o_h2;

	int m = M;
	int n = N;
	int k = K;

	int r,c;

	cl_int err;

	int size = 16;
	size_t local[2] = {size, size};
	size_t global[2] = {shrRoundUp(local[0],M) , shrRoundUp(local[1],N)};

	cl_platform_id platform;
	cl_device_id device;
	cl_context context;
	cl_command_queue command;
	cl_program program;
	cl_kernel kernel;

	cl_event event = NULL;

	m1_h = (float*)malloc(M*K * sizeof(float));
	m2_h = (float*)malloc(K*N * sizeof(float));
	o_h = (float*)malloc(M*N * sizeof(float));
	o_h2 = (float*)malloc(M*N * sizeof(float));

	// Initialize host array and copy it to CUDA device
	for (r = 0 ; r < M; ++r) {
		for (c = 0 ; c < K; ++c) {
			m1_h[r*K + c] = rand()/(float)RAND_MAX;
		}
	}

	for (r = 0 ; r < K; ++r) {
		for (c = 0 ; c < N; ++c) {
			m2_h[r*N + c] = rand()/(float)RAND_MAX;
		}
	}

	//get platform id
	err = clGetPlatformIDs(1, &platform, NULL);
	if(err!= CL_SUCCESS){
		printf( "clGetPlatformIDs() failed with %d\n", err );
		exit(-1);
	}

	//get device id
	err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
	if(err!= CL_SUCCESS){
		printf( "clGetDeviceIDs() failed with %d\n", err );
		exit(-1);
	}

	//Create Context
	context = clCreateContext(0, 1, &device, NULL, NULL, &err);
	if(err != CL_SUCCESS){
		printf( "clCreateContext() failed with %d\n", err );
		exit(-1);
	}

	//create command queue
	command = clCreateCommandQueue(context, device, 0, &err);
	if(err != CL_SUCCESS){
		printf( "clCreateCommandQueue() failed with %d\n", err );
		exit(-1);
	}

	// Setup clAmdBlas
	err = clAmdBlasSetup();
	if(err != CL_SUCCESS){
		printf("clAmdBlasSetup() failed with %d\n",err);
		clReleaseCommandQueue(command);
		clReleaseContext(context);
		exit(-1);
	}	

	//Create program
	program = clCreateProgramWithSource(context, 1, (const char **)& KernelSource3, NULL, &err);
	if(err != CL_SUCCESS){
		printf( "clCreateProgramWithSource() failed with %d\n", err );
		exit(-1);
	}

	//build program
	err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (err != CL_SUCCESS) {
		printf( "clBuildProgram() failed with %d\n", err );
		if(err == CL_BUILD_PROGRAM_FAILURE){
			size_t log_size;
			char *log;
			clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);

			log = (char *)malloc(log_size);

			clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);

			printf("%s\n", log);

			exit(-1);
		}
		exit(-1);
	}


	//create kernel
	kernel = clCreateKernel(program, "kernel3", &err);
	if (err != CL_SUCCESS) {
		printf( "clBuildKernel() failed with %d\n", err );
		exit(-1);
	}


	//create buffer
	m1_d = clCreateBuffer(context, CL_MEM_READ_ONLY, M * K * sizeof(float),NULL,&err);
	if (err != CL_SUCCESS) {
		printf("Can't Create buffer a! %d\n",err);
		exit(-1);
	}
	m2_d = clCreateBuffer(context, CL_MEM_READ_ONLY, K * N * sizeof(float),NULL,&err);
	if (err != CL_SUCCESS) {
		printf("Can't Create buffer b! %d\n",err);
		exit(-1);
	}
	o_d = clCreateBuffer(context, CL_MEM_READ_WRITE, M * N * sizeof(float),NULL,&err);
	if (err != CL_SUCCESS) {
		printf("Can't Create buffer c! %d\n",err);
		exit(-1);
	}

	//enqueue wrte buffer
	err = clEnqueueWriteBuffer(command, m1_d, CL_TRUE, 0, M * K * sizeof(float),m1_h,0 ,NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write Buffer A! %d\n",err);
		exit(-1);
	}
	err = clEnqueueWriteBuffer(command, m2_d, CL_TRUE, 0, K * N * sizeof(float),m2_h,0 ,NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write Buffer B! %d\n",err);
		exit(-1);
	}
	err = clEnqueueWriteBuffer(command, o_d, CL_TRUE, 0, M * N * sizeof(float),o_h,0 ,NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write Buffer C! %d\n",err);
		exit(-1);
	}


	//set function arguments

	err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &o_d);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment c!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &m1_d);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment a!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &m2_d);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment b!\n");
		exit(-1);
	}
	err = clSetKernelArg(kernel, 3, sizeof(int), &m);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment a!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 4, sizeof(int), &n);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment b!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 5, sizeof(int), &k);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment c!\n");
		exit(-1);
	}



	gettimeofday(&time_start, NULL);

	err = clEnqueueNDRangeKernel(command, kernel, 2, NULL,global,local, 0, NULL, &event);

	err |= clWaitForEvents(1, &event);

	gettimeofday(&time_end, NULL);

	if (err != CL_SUCCESS) {
		printf("Can't Enqueue NDRange Kernel %d!\n",err);
		exit(-1);
	}


	err = clFinish(command);
	if (err != CL_SUCCESS) {
		printf("Can't finish!\n");
		exit(-1);
	}

	timersub(&time_end, &time_start, &time_res);

	printf("%ld,%03ld ms\n", time_res.tv_sec, time_res.tv_usec);

	err = clEnqueueReadBuffer( command, o_d, CL_TRUE, 0,N * N * sizeof(float) , o_h , 0, NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Read Buffer\n");
		exit(-1);
	}

	err = clAmdBlasSgemm(order, transA, transB, M, N, K, alpha, m1_d, K, m2_d,
			N, beta, o_d, N, 1, &command, 0, NULL, &event);
	if(err != CL_SUCCESS){
		printf("clAmdBlasSgemm() failed with %d\n", err);
		exit(-1);
	}
	else{
		err = clWaitForEvents(1, &event);
		err = clEnqueueReadBuffer(command, o_d, CL_TRUE, 0,
				M * N * sizeof(float),
				o_h2, 0, NULL, NULL);
	}

	int i,j;
	for(i = 0; i < N ; i++){
		for(j = 0 ; j < N ; j++){
			if( (o_h[i*N+j] - o_h2[i*N+j]) > 0.001 || (o_h[i*N+j] - o_h2[i*N+j]) < -0.001){
				count++;
			}
		}
	}

	printf("number of differences : %d\n",count);


	clReleaseMemObject(m1_d);
	clReleaseMemObject(m2_d);
	clReleaseMemObject(o_d);

	clReleaseCommandQueue(command);
	clReleaseContext(context);

	return 0;
}
